//----------------------------------------------------------------------------
// M A R I T I M E  R E S E A R C H  I N S T I T U T E  N E T H E R L A N D S
//----------------------------------------------------------------------------
// Copyright (C) 2011 - MARIN - All rights reserved - http://www.marin.nl
//----------------------------------------------------------------------------
// Program    : mermaid
// Module     : vbm
// File       : defs.h
// Author     : M.A. de Jong
//----------------------------------------------------------------------------

#ifndef DEFS_H
#define	DEFS_H

// ===========================================================================
// "Adjustable parameters"
// ===========================================================================

#define DIM_COMPUTE_BLOCK        32
// About:/ This number must be set beforehand! For older devices (compute
// capability < 2.0) set this number to 16. For new devices (Fermi, compute
// capability 2.0 or 2.1) set this number to 32 for best performance (however,
// 16 is of course also allowed). [Default: 16]


#define BORDER_WIDTH             16
// About:/ This number represents the width of the borders that are used to
// make sure that global memory reads are as much as possible performed in a 
// coalesced manner (which yields optimal throughput). A good choice is 
// BORDER_WIDTH = 16. [Default: 16]


#define USE_PAGELOCKED           0   
// About:/ When set to 1 the bandwidth of device to host transfers may increase 
// somewhat (on older devices). It is used to compute the dot products.
// [Default: 0]


#define MAXITER                  99 
// About:/ The maximal number of CG iterations allowed. For a real-time solver
// the CG algorithm should convergence fast enough, when the number of CG
// iterations in the RRB-solver it too large, something is wrong, possibly the
// RRB-solver fails to solve the particular problem Ax = b. [Default: 99]


#define NUM_THREADS              256
// About:/ This number is the number of threads in 1D thread-blocks in "not so
// important kernels"; 64, 128, 256, 512 are almost equally good. 
// [Default: 256]


#define USE_TEXTURES             1
// About:/ When set to 1, some kernels will use 2D textures to retrieve data
// from the global memory, on older architectures yielding significantly higher 
// throughput due to the fact that fetches are cached. However, on new 
// architectures L1 cache is available which makes that reading with a shift
// comes without a penalty, so textures will not do much. [Default: 1]
// ------
// Extra note: in case of doubles USE_TEXTURES will be set to 0 automatically
// via preprocessor directives below


// ===========================================================================
// "Essential parameters and stuff". Do not change.
// ===========================================================================
#define Bx                       blockDim.x
#define bx                       blockIdx.x
#define tx                       threadIdx.x
#define By                       blockDim.y
#define by                       blockIdx.y
#define ty                       threadIdx.y
#define Gx                       gridDim.x
#define Gy                       gridDim.y

#define BORDER_WIDTH2            (2 * BORDER_WIDTH)

#define T1TO4_BF                 (DIM_COMPUTE_BLOCK / 4)
// About:/ Block-factor to go from 1 to 4 grids and vice versa

#define R2TO4_BF                 1
// About:/ Block-factor to go from r2 nodes to a new grid with the 
// r1/r2/b1/b2 storage format and vice versa

#define DOTP_BF                  (DIM_COMPUTE_BLOCK / 2)
// About:/ Block-factor used in the dot product kernel

#define PREC_K1                  1
#define PREC_K2                  2
#define PREC_K3                  3
#define PREC_K4                  4
#define PREC_K5                  5
#define SUBS_K1                  6
#define SUBS_K2                  7
#define SOLV_K1                  8
#define SOLV_K2                  9
#define SOLV_K3                  10
#define SOLV_K4                  11
#define SOLV_K5                  12
#define SOLV_K6                  13
#define MATV_K1                  14
#define MATV_K2                  15
#define AXPY_K1                  16

#ifdef USEDOUBLE
  #undef  DIM_COMPUTE_BLOCK
  #define DIM_COMPUTE_BLOCK      16
  #undef  USE_TEXTURES
  #define USE_TEXTURES           false
#endif

#define NOCLOCK                  false
#define CLOCK                    true
#define NOINIT                   false
#define INIT                     true

#if USE_TEXTURES
  #define __FETCH_CC(i, j, ld)   tex2D(texCC, i, j)
  #define __FETCH_NN(i, j, ld)   tex2D(texNN, i, j)
  #define __FETCH_EE(i, j, ld)   tex2D(texEE, i, j)
  #define __FETCH_SS(i, j, ld)   tex2D(texSS, i, j)
  #define __FETCH_WW(i, j, ld)   tex2D(texWW, i, j)
  #define __FETCH_NE(i, j, ld)   tex2D(texNE, i, j)
  #define __FETCH_SE(i, j, ld)   tex2D(texSE, i, j)
  #define __FETCH_SW(i, j, ld)   tex2D(texSW, i, j)
  #define __FETCH_NW(i, j, ld)   tex2D(texNW, i, j)
  #define __FETCH_XX(i, j, ld)   tex2D(texRef, i, j)
  #define __FETCH_YY(i, j, ld)   tex2D(texRef, i, j)
#else
  #define __FETCH_CC(i, j, ld)   g.cc[(j) * ld + (i)]
  #define __FETCH_NN(i, j, ld)   g.nn[(j) * ld + (i)]
  #define __FETCH_EE(i, j, ld)   g.ee[(j) * ld + (i)]
  #define __FETCH_SS(i, j, ld)   g.ss[(j) * ld + (i)]
  #define __FETCH_WW(i, j, ld)   g.ww[(j) * ld + (i)]
  #define __FETCH_NE(i, j, ld)   g.ne[(j) * ld + (i)]
  #define __FETCH_SE(i, j, ld)   g.se[(j) * ld + (i)]
  #define __FETCH_SW(i, j, ld)   g.sw[(j) * ld + (i)]
  #define __FETCH_NW(i, j, ld)   g.nw[(j) * ld + (i)]
  #define __FETCH_XX(i, j, ld)   x[(j) * ld + (i)]
  #define __FETCH_YY(i, j, ld)   y[(j) * ld + (i)]
#endif

// Default thread-block dimensions 
#define NUM_THREADS_X            DIM_COMPUTE_BLOCK
#define NUM_THREADS_Y           (DIM_COMPUTE_BLOCK / 4)


#endif // DEFS_H
